// 包含标准输入输出库，提供printf函数
#include <stdio.h>    
// 包含CUDA运行时API头文件，提供GPU编程接口
#include <cuda.h>     

//----------------------------------------------------------
// 核函数1：打印三维块索引 + 线程三维索引 + 线程线性索引
//----------------------------------------------------------
__global__ void print_idx_kernel() {
    // 将三维线程索引转换为线性索引（z*y*x + y*x + x）
    int index = threadIdx.z * blockDim.x * blockDim.y  // z维度偏移量（z方向每层有x*y个线程）
              + threadIdx.y * blockDim.x              // y维度偏移量（y方向每行有x个线程）
              + threadIdx.x;                          // x维度基础索引

    // 格式化输出：
    // blockIdx.z/y/x表示块在网格中的三维坐标（注意此处是z,y,x的反向顺序）
    // index为计算得到的线性索引
    // threadIdx.z/y/x表示线程在线程块内的三维坐标
    printf("block idx:(%3d,%3d,%3d), thread idx:%3d,thread idx:(%3d,%3d,%3d)\n",
           blockIdx.z, blockIdx.y, blockIdx.x,   // 输出块的三维索引（z,y,x顺序）
           index,                                // 线程在线程块内的线性索引
           threadIdx.z, threadIdx.y, threadIdx.x); // 线程的三维索引（z,y,x顺序）

    }

//----------------------------------------------------------
// 核函数2：打印网格和线程块的维度（已定义但未调用）
//----------------------------------------------------------
__global__ void print_dim_kernel() {
    // 输出网格和线程块的三维尺寸
    // 注意CUDA维度顺序为x,y,z，这里使用反向顺序z,y,x输出
    printf("grid dimension:(%3d,%3d,%3d), block dimension:(%3d,%3d,%3d)\n",
        gridDim.z, gridDim.y, gridDim.x,   // 网格尺寸（z,y,x方向块数量）
        blockDim.z, blockDim.y, blockDim.x); // 线程块尺寸（z,y,x方向线程数）
}

//----------------------------------------------------------
// 核函数3：打印全局线程索引
//----------------------------------------------------------
__global__ void print_thread_idx_kernel()
{   
    // 计算线程块内总线程数（z*y*x）
    int bsize = blockDim.z * blockDim.y * blockDim.x;

    // 计算当前线程块的全局线性索引：
    int bIndex = blockIdx.z * gridDim.x * gridDim.y +  // z维度偏移（每层有x*y个块）
                 blockIdx.y * gridDim.x +              // y维度偏移（每行有x个块）
                 blockIdx.x;                           // x维度基础索引

    // 计算线程在线程块内的线性索引（同核函数1）
    int tIndex = threadIdx.z * blockDim.x * blockDim.y +
                 threadIdx.y * blockDim.x +
                 threadIdx.x;
   
    // 计算全局线程索引：块索引*块大小 + 线程块内索引
    int index = bIndex * bsize + tIndex;
    
    // 输出块索引、线程块内索引、全局线程索引
    printf("block idx:%3d, thread idx in block: %3d, thread idx: %3d\n",
           bIndex, tIndex, index);
}

//----------------------------------------------------------
// 核函数4：打印二维坐标
//----------------------------------------------------------
__global__ void print_cord_kernel()
{   
    // 计算线程在线程块内的线性索引
    int index = threadIdx.z * blockDim.x * blockDim.y +
                threadIdx.y * blockDim.x +
                threadIdx.x;
    
    // 计算全局二维坐标：
    int x = blockIdx.x * blockDim.x + threadIdx.x; // 全局x坐标 = 块x偏移 + 线程x索引
    int y = blockIdx.y * blockDim.y + threadIdx.y; // 全局y坐标 = 块y偏移 + 线程y索引
    
    // 输出块三维索引、线程块内索引、全局二维坐标
    // 注意z维度未参与坐标计算（适用于二维情况）
    printf("block idx:(%3d,%3d,%3d), thread idx: %3d, cord:(%3d, %3d)\n", 
       blockIdx.z, blockIdx.y, blockIdx.x, // 块的三维索引
       index,                              // 线程块内线性索引
       x, y);                              // 全局二维坐标
}

//----------------------------------------------------------
// 主机端配置函数
//----------------------------------------------------------
void helloFromGPU() {
    int inputwidth = 4;        // 假设数据宽度为8（用于计算网格尺寸）
    int block_dim = 2;         // 每个维度的线程块大小（创建2x2线程块）
    int grid_dim = inputwidth / block_dim; // 计算网格维度（8/2=4）
    
    // 定义线程块结构（二维：2x2x1）
    dim3 block(block_dim, block_dim); // 等效dim3(2,2,1)
    // 定义网格结构（二维：4x4x1）
    dim3 grid(grid_dim, grid_dim);    // 等效dim3(4,4,1)
    
    // 启动核函数（当前启用了print_thread_idx_kernel）
    
    print_thread_idx_device(grid, block);
}

void print_thread_idx_device(dim3 grid, dim3 block)
{   
    print_thread_idx_kernel<<<grid, block>>>();
}

//----------------------------------------------------------
// 主函数
//----------------------------------------------------------
int main() {
    helloFromGPU();  // 调用GPU函数

    // CUDA错误检查机制
    cudaError_t launchErr = cudaGetLastError();  // 检查核函数启动错误
    if (launchErr != cudaSuccess) {
        printf("Kernel launch error: %s\n", cudaGetErrorString(launchErr));
        return -1;
    }

    cudaError_t syncErr = cudaDeviceSynchronize(); // 同步设备并检查运行时错误
    if (syncErr != cudaSuccess) {
        printf("Device sync error: %s\n", cudaGetErrorString(syncErr));
        return -2;
    }

    printf("Hello, World! from CPU#######\n");  // CPU端输出
    return 0;
}